//
//  DeferredLightingView.m
//  DeferredLighting
//
//  Created by Holmes Futrell on 2/8/11.
//  Copyright 2011 __MyCompanyName__. All rights reserved.
//

#import "DeferredLightingView.h"
#import "clb_error.h"
#import "clb_image.h"
#import "clb_event_profiling.h"

enum {
	kPassGBuffer,
	kPassShade
};

enum {
	kFinalBlitModeColor=0,
	kFinalBlitModeDepth,
	kFinalBlitModeDepthPrime,
	kFinalBlitModeNormal,
	kFinalBlitModeDiffuse,
	kFinalBlitModeSpecular,
	kFinalBlitModeDebug1,
	kFinalBlitModeDebug2,
	kNumFinalBlitModes
};

@implementation DeferredLightingView

- (id)initWithFrame:(NSRect)frameRect pixelFormat:(NSOpenGLPixelFormat *)format {

	if ( self = [super initWithFrame: frameRect pixelFormat: format] ) {
		fbo = 0;
		colorTextureID = 0;
		depthTextureID = 0;
		normalTextureID = 0;
		lightbufferDiffuseTextureID = 0;
		lightbufferSpecularTextureID = 0;
		return self;
	}
	else {
		return nil;
	}
	
}

- (BOOL)loadOpenCL {
	
	simpleCompute = [[SimpleCompute alloc] initForGPU];
	
	block_size[0] = 8;
	block_size[1] = 8;
	tile_size[0] = 8;
	tile_size[1] = 8;
	threadsPerTile = 64;
	
	int pixelsPerTile[2];
	pixelsPerTile[0] = block_size[0] * tile_size[0];
	pixelsPerTile[1] = block_size[1] * tile_size[1];

	NSString *options = [NSString stringWithFormat:@"-DBLOCK_X=%d -DBLOCK_Y=%d -DPIXELS_X_PER_TILE=%d -DPIXELS_Y_PER_TILE=%d",
						 block_size[0], block_size[1],
						 pixelsPerTile[0], pixelsPerTile[1]];	
	NSString *optionsBinning = [NSString stringWithFormat:@"-DBLOCK_X=%d -DBLOCK_Y=%d -DTHREADS_PER_TILE=%d -DTILE_X=%d -DTILE_Y=%d",\
								block_size[0], block_size[1], \
								threadsPerTile,\
								tile_size[0], tile_size[1]];	

	computeProgram = [[ComputeProgram alloc] initWithContext: simpleCompute.context
													deviceID: simpleCompute.deviceID
														file: [[NSBundle mainBundle] pathForResource:@"simpleGenLightbuffer" ofType:@"cl"]
												  kernelName: @"simpleGenLightBuffer"
												 buildOptions: options];
	
	
	perBlockAttributesProgram = [[ComputeProgram alloc] initWithContext: simpleCompute.context
															   deviceID: simpleCompute.deviceID
																   file: [[NSBundle mainBundle] pathForResource:@"perblockattributes" ofType:@"cl"]
															 kernelName: @"perblockattributes"
														   buildOptions: options];
	
	binningProgram = [[ComputeProgram alloc] initWithContext: simpleCompute.context
															   deviceID: simpleCompute.deviceID
																   file: [[NSBundle mainBundle] pathForResource:@"binning" ofType:@"cl"]
															 kernelName: @"binLights"
														   buildOptions: optionsBinning];
	
	int numParticles =  1024;
	
	cl_int error;
	lightBins = clCreateBuffer(simpleCompute.context,
									   CL_MEM_READ_WRITE,
									   sizeof(int) * numParticles * 32 * 32,
									   NULL,
									   &error);
	clb_check_error_simple(error);
	
	binCounts = clCreateBuffer(simpleCompute.context,
							   CL_MEM_READ_WRITE,
							   sizeof(int) * numParticles * 32 * 32,
							   NULL,
							   &error);
	clb_check_error_simple(error);
	
	
	//num_per_block_attributes = 10;
	
	//int i;
	//for (i=0; i<10; i++) {
		//cl_int err;
		//size_t size_bytes = num_per_block_attributes * ( max_resolution / block_size[0] ) * ( max_resolution / block_size[1] ) * sizeof(float); 
		//per_block_attributes = clCreateBuffer( simpleCompute.context, CL_MEM_READ_WRITE, size_bytes, NULL, &err);
		//clb_check_error(err, "allocating per block attribute");
	//}
	
	nearestSampler = clCreateSampler(simpleCompute.context, CL_FALSE, CL_ADDRESS_CLAMP, CL_FILTER_NEAREST, &error);
	clb_check_error_simple(error);
	
	//clb_print_supported_2D_image_formats(simpleCompute.context, CL_MEM_READ_WRITE);
	//clb_print_supported_2D_image_formats(simpleCompute.context, CL_MEM_READ_ONLY);
	//clb_print_supported_2D_image_formats(simpleCompute.context, CL_MEM_WRITE_ONLY);

	particleSystem = [[ParticleSystem alloc] initWithNumParticles: numParticles simpleCompute: simpleCompute];
	
	
	return YES;
	
}

- (BOOL)loadShaders
{
		
	tangentShader = [[ShaderProgram alloc] initWithVertexShader: [[NSBundle mainBundle] pathForResource:@"tangentspace" ofType:@"vsh"]
												 fragmentShader: [[NSBundle mainBundle] pathForResource:@"tangentspace" ofType:@"fsh"]];
	
	[tangentShader linkProgram];
	
	
	basicShader = [[ShaderProgram alloc] initWithVertexShader: [[NSBundle mainBundle] pathForResource:@"basic" ofType:@"vsh"]
											   fragmentShader: [[NSBundle mainBundle] pathForResource:@"basic" ofType:@"fsh"]];
	[basicShader linkProgram];
	
	blitShader = [[ShaderProgram alloc] initWithVertexShader: [[NSBundle mainBundle] pathForResource:@"blit" ofType:@"vsh"]
											   fragmentShader: [[NSBundle mainBundle] pathForResource:@"blit" ofType:@"fsh"]];
	[blitShader linkProgram];
	
	lightBufferShader = [[ShaderProgram alloc] initWithVertexShader: [[NSBundle mainBundle] pathForResource:@"lightbuffer" ofType:@"vsh"]
											  fragmentShader: [[NSBundle mainBundle] pathForResource:@"lightbuffer" ofType:@"fsh"]];
	
	[lightBufferShader linkProgram];
	
	
	visualizeColor = [[ShaderProgram alloc] initWithVertexShader: [[NSBundle mainBundle] pathForResource:@"visualize" ofType:@"vsh"]
						fragmentShader: [[NSBundle mainBundle] pathForResource:@"visualize-color" ofType:@"fsh"]];
	
	[visualizeColor linkProgram];
	
	simpleGenLightbufferShader = [[ShaderProgram alloc] initWithVertexShader: [[NSBundle mainBundle] pathForResource:@"simpleGenLightbuffer" ofType:@"vsh"]
													 fragmentShader: [[NSBundle mainBundle] pathForResource:@"simpleGenLightbuffer" ofType:@"fsh"]];
	
	[simpleGenLightbufferShader linkProgram];
	
    // Get uniform locations.
	
	count = 0;
	fps = 0;
	[NSTimer scheduledTimerWithTimeInterval: 1 target: self selector:@selector(countFrames:) userInfo: nil repeats: YES];
	[NSTimer scheduledTimerWithTimeInterval: 0 target: self selector:@selector(animationTimer:) userInfo: nil repeats: YES];

	rb_check_error_simple();
	
    return TRUE;
}

-(void)animationTimer:(NSTimer *)timer {
	//rot += 0.001;
	[self setNeedsDisplay: YES];
}

-(void)countFrames:(NSTimer *)timer {
	fps = count;
	count = 0;
}

-(void)visualizeColorWithImageSize:(CGSize)size dstRect:(CGRect)dst {
	
	[visualizeColor useProgram];
	
	mat4 projection = rb_mat4_ortho(0, viewportWidth, 0, viewportHeight, -10, 10);
	[visualizeColor useProgram];
	
	glUniformMatrix4fv([visualizeColor locationForUniform:@"projectionMatrix"], 1, 0, projection.entries);
	
	int i1 = [visualizeColor getLocationForAttribute:@"position_in"];
	int i2 = [visualizeColor getLocationForAttribute:@"uv_in"];

	glUniform1i([visualizeColor locationForUniform:@"sampler"], 0);
	
	rb_check_error_simple();
	
	float xMin = dst.origin.x;
	float xMax = dst.origin.x + dst.size.width;
	float yMin = dst.origin.y;
	float yMax = dst.origin.y + dst.size.height;

	GLfloat texcoords[] = { 0.0, 0.0, 1, 0.0, 1, 1, 1, 1, 0.0, 1, 0.0, 0.0 };
	GLfloat coords[] = { xMin, yMin, xMax, yMin, xMax, yMax, xMax, yMax, xMin, yMax, xMin, yMin };

	glEnableVertexAttribArray(i1);
	glVertexAttribPointer(i1, 2, GL_FLOAT, GL_FALSE, 0, coords); 
	rb_check_error_simple();
	glEnableVertexAttribArray(i2);
	glVertexAttribPointer(i2, 2, GL_FLOAT, GL_FALSE, 0, texcoords); 
	rb_check_error_simple();

	glDrawArrays(GL_TRIANGLES, 0, 6);
	rb_check_error_simple();
	
	glDisableVertexAttribArray(i1);
	glDisableVertexAttribArray(i2);
	rb_check_error_simple();
	
	
}


-(void)drawSquare:(int)positionAttributeLocation {
	
	rb_check_error_simple();
	GLfloat coords[] = { 0.0, 0.0, 1.0, 0.0, 1.0, 1.0, 1.0, 1.0, 0.0, 1.0, 0.0, 0.0 };
	glEnableVertexAttribArray(positionAttributeLocation);
	rb_check_error_simple();
	glVertexAttribPointer(positionAttributeLocation, 2, GL_FLOAT, GL_FALSE, 0, coords); 
	rb_check_error_simple();
	glDrawArrays(GL_TRIANGLES, 0, 6);
	rb_check_error_simple();	
	glDisableVertexAttribArray(positionAttributeLocation);
	
	rb_check_error_simple();

	
}

-(mat4)getModelview {

	mat4 modelView = rb_mat4_rotate( -3.14159f / 2.0f, 1, 0, 0 );
	modelView = rb_mat4_mul(rb_mat4_rotate( M_PI, 1, 0, 0 ), modelView);
	modelView = rb_mat4_mul(rb_mat4_rotate( rot, 0, 1, 0 ), modelView);
	
	modelView = rb_mat4_mul(rb_mat4_rotate( 30.0f * M_PI / 180.0f, 1, 0, 0 ), modelView);

	
	modelView = rb_mat4_mul(rb_mat4_translate(make_vec3(0,-0.3,-3.0)), modelView);

	return modelView;
	
}

-(void)drawPass:(int)pass {
			
	mat4 modelView = [self getModelview];
	mat4 projection = rb_mat4_frustum( left, right, bottom, top, zNear, zFar);
	
	//mat4 projection = rb_mat4_perspective( fovyInDegrees, aspect, zNear, zFar);
	
	mat4 modelViewProjection = rb_mat4_mul(projection, modelView);	
			
	if ( pass == kPassGBuffer ) {
						
		glBindFramebuffer(GL_FRAMEBUFFER, fbo);
		
		glClearColor(0.5, 0.5, 0.0, 0.0);		
		glFramebufferTexture2D(GL_FRAMEBUFFER, GL_COLOR_ATTACHMENT0, GL_TEXTURE_2D, normalTextureID, 0);
		glFramebufferTexture2D(GL_FRAMEBUFFER, GL_DEPTH_ATTACHMENT, GL_TEXTURE_2D, depthTextureID, 0);

		GLuint status = glCheckFramebufferStatus(GL_FRAMEBUFFER);
		if ( status != GL_FRAMEBUFFER_COMPLETE ) {
			NSLog(@"Incomplete framebuffer %s (pass = %s)", rbunny_framebuffer_status_to_string(status), "kPassGBuffer" );
		}				
		
		// Use shader program.
		[tangentShader useProgram];
		
		glUniformMatrix4fv( [tangentShader locationForUniform:@"modelviewMatrix"], 1, 0, modelView.entries);
		glUniformMatrix4fv( [tangentShader locationForUniform:@"modelviewProjectionMatrix"], 1, 0, modelViewProjection.entries);		
		glUniform1f([tangentShader locationForUniform:@"specularPower"],	75.0f);
		glUniform1i([tangentShader locationForUniform:@"normalsSampler"],	0);
		
		glActiveTexture(GL_TEXTURE0);
		glBindTexture(GL_TEXTURE_2D, normals.descriptor);
				
		glEnable(GL_DEPTH_TEST);	
		glDepthFunc(GL_LEQUAL);

		glDepthMask(GL_TRUE);
		glClear(GL_COLOR_BUFFER_BIT | GL_DEPTH_BUFFER_BIT);
		
		[model1 bindVertexArray];
		[model1 draw];	
		[model1 unbindVertexArray];
				
		glActiveTexture(GL_TEXTURE0);
		glBindTexture(GL_TEXTURE_2D, 0);

		
	}
	else if ( pass == kPassShade ) {
		
		glClearColor(0.0, 0.0, 0.0, 1.0);		
		glBindFramebuffer(GL_FRAMEBUFFER, fbo);
		glFramebufferTexture2D(GL_FRAMEBUFFER, GL_COLOR_ATTACHMENT0, GL_TEXTURE_2D, colorTextureID, 0);
		//glFramebufferTexture2D(GL_FRAMEBUFFER, GL_DEPTH_ATTACHMENT, GL_TEXTURE_2D, depthTextureID, 0);
		
		GLuint status = glCheckFramebufferStatus(GL_FRAMEBUFFER);
		if ( status != GL_FRAMEBUFFER_COMPLETE ) {
			NSLog(@"Incomplete framebuffer %s (pass = %s)", rbunny_framebuffer_status_to_string(status), "kPassGBuffer" );
		}				
		
		// Use shader program.
		[lightBufferShader useProgram];
		rb_check_error_simple();

		glUniformMatrix4fv( [lightBufferShader locationForUniform:@"modelviewMatrix"], 1, 0, modelView.entries);
		glUniformMatrix4fv( [lightBufferShader locationForUniform:@"modelviewProjectionMatrix"], 1, 0, modelViewProjection.entries);		
		glUniform1i([lightBufferShader locationForUniform:@"textureSampler"],	0);
		glUniform1i([lightBufferShader locationForUniform:@"lightbufferDiffuseSampler"],	1);
		glUniform1i([lightBufferShader locationForUniform:@"lightbufferSpecularSampler"],	2);
		glUniform2f([lightBufferShader locationForUniform:@"imageSize"], viewportWidth, viewportHeight);

		rb_check_error_simple();

		glActiveTexture(GL_TEXTURE0);
		glBindTexture(GL_TEXTURE_2D, texture.descriptor);
		rb_check_error_simple();

		glActiveTexture(GL_TEXTURE1);
		glBindTexture(GL_TEXTURE_2D, lightbufferDiffuseTextureID); // change to diffuse light
		rb_check_error_simple();

		glActiveTexture(GL_TEXTURE2);
		glBindTexture(GL_TEXTURE_2D, lightbufferSpecularTextureID); // change to specular light
		rb_check_error_simple();

		glEnable(GL_DEPTH_TEST);
		glDepthMask(GL_FALSE);
		glClear(GL_COLOR_BUFFER_BIT);
		rb_check_error_simple();

		[model2 bindVertexArray];
		[model2 draw];	
		[model2 unbindVertexArray];
		rb_check_error_simple();
		
		[particleSystem drawParticlesWithModelviewMatrix: modelView
										projectionMatrix: projection];
		
		glActiveTexture(GL_TEXTURE0);
		glBindTexture(GL_TEXTURE_2D, 0);
		glActiveTexture(GL_TEXTURE1);
		glBindTexture(GL_TEXTURE_2D, 0);
		glActiveTexture(GL_TEXTURE2);
		glBindTexture(GL_TEXTURE_2D, 0);
		rb_check_error_simple();

		
	}
	
	rb_check_error_simple();
			
	glBindFramebuffer(GL_FRAMEBUFFER, 0);

	
}

-(cl_mem)createImageForOpenCLOnlyWithSize:(CGSize)size format:(cl_image_format)format {

	cl_int err = CL_SUCCESS;
	

	cl_mem result = clCreateImage2D(simpleCompute.context, CL_MEM_READ_WRITE, &format, size.width, size.height, 0, NULL, &err);
	clb_check_error(err, "clCreateImage2D");
	
	return result;
	
}

-(cl_mem)createComputeResultFromTexture:(GLuint)textureID mode:(cl_mem_flags)flags
{
		
	int err = 0;
    
	cl_mem result = clCreateFromGLTexture2D(simpleCompute.context, 
						flags,
						GL_TEXTURE_2D,
						0,
						textureID,
						&err);
	
	clb_check_error(err, "clCreateFromGLTexture2D");
	
    if (!result)
    {
        clbfatalError("Failed to create OpenGL texture reference!");
		return 0;
    }
	
    return result;
}


-(void)generateLightBufferOpenCL {
	
	cl_command_queue queue = simpleCompute.queue;
	//cl_context context = simpleCompute.context;	
	
	cl_int error = CL_SUCCESS;
	
	int imageSize[2];
	imageSize[0] = viewportWidth;
	imageSize[1] = viewportHeight;	
	
	// allow OpenCL to start operating on these objects (since they were created by OpenGL)
	error |= clEnqueueAcquireGLObjects(queue, 1, &diffuseImage, 0, 0, 0);
	error |= clEnqueueAcquireGLObjects(queue, 1, &specularImage, 0, 0, 0);
	error |= clEnqueueAcquireGLObjects(queue, 1, &normSpecImage, 0, 0, 0);
	error |= clEnqueueAcquireGLObjects(queue, 1, &depthImage, 0, 0, 0);
	clb_check_error_simple(error);
	
	cl_int work_dim = 2;
	
	//printf("image size = %d, %d\n", imageSize[0], imageSize[1]);
	//printf("num blocks = %d, %d\n", num_blocks[0], num_blocks[1]);
	
	cl_event kernelFinish1Event;
	cl_event kernelFinish2Event;
	cl_event kernelFinish3Event;

	{
		
#if 1
		size_t local_work_size[work_dim];
		local_work_size[0] = 8;
		local_work_size[1] = 8;

		size_t global_work_size[work_dim];
		global_work_size[0] = ((num_blocks[0] + local_work_size[0] - 1) / local_work_size[0]) * local_work_size[0];
		global_work_size[1] = ((num_blocks[1] + local_work_size[1] - 1) / local_work_size[1]) * local_work_size[1];
#endif 
		
#if 0
		size_t local_work_size[work_dim];
		local_work_size[0] = block_size[0];
		local_work_size[1] = block_size[1];
		size_t global_work_size[work_dim];
		global_work_size[0] = num_blocks[0] * block_size[0];
		global_work_size[1] = num_blocks[1] * block_size[1];
#endif 
		
		//printf("per block: global work size = %d, %d.  Local work size = %d, %d\n",
		//	   global_work_size[0], global_work_size[1], local_work_size[0], local_work_size[1]);
		
		cl_kernel per_block_kernel = perBlockAttributesProgram.kernel;
		clb_check_error_simple(clSetKernelArg(per_block_kernel, 0, sizeof(cl_mem), &normSpecImage));
		clb_check_error_simple(clSetKernelArg(per_block_kernel, 1, sizeof(cl_mem), &depthImage));
		clb_check_error_simple(clSetKernelArg(per_block_kernel, 2, sizeof(cl_sampler), &nearestSampler));
		clb_check_error_simple(clSetKernelArg(per_block_kernel, 3, 2 * sizeof(int), &imageSize));
		clb_check_error_simple(clSetKernelArg(per_block_kernel, 4, 2 * sizeof(int), &num_blocks));
		clb_check_error_simple(clSetKernelArg(per_block_kernel, 5, sizeof(cl_mem), &per_block_image1));
		clb_check_error_simple(clSetKernelArg(per_block_kernel, 6, sizeof(cl_mem), &per_block_image2));
		error = clEnqueueNDRangeKernel(queue, per_block_kernel, work_dim, NULL, global_work_size, local_work_size, 0, NULL, &kernelFinish1Event);
		clb_check_error_simple(error);
		
	}

	float		Q1 = 2.0f * zNear / ( right - left );
	float		Q2 = 2.0f * zNear / ( top - bottom );
	float		C  = - (zFar + zNear) / (zFar - zNear);
	float		D  = - (2.0f * zFar * zNear) / (zFar - zNear);	
	
	cl_mem lights = [particleSystem eyespaceParticles];
	int numLights = particleSystem.numParticles;

	int numTiles[2];
	numTiles[0] = (num_blocks[0]+tile_size[0]-1) / tile_size[0];
	numTiles[1] = (num_blocks[1]+tile_size[1]-1) / tile_size[1];
	
	
	{
	/*
	 __read_only	 image2d_t input_image,
	 const int2 numTiles,
	 __global float4 *lights,
	 const int numLights,
	 __global float4 *lightBins,
	 const		float		zNear,
	 const		float		Q1, // 2 * zNear / ( right - left )
	 const		float		Q2, // 2 * zNear / ( top - bottom )
	 const		float		C,  // - (zFar + zNear) / (zFar - zNear)
	 const		float		D,  // - (2 zFar zNear) / (zFar - zNear)
	 sampler_t	nearestSampler,
	 int2 imageSize
	 */
		cl_kernel binningKernel = binningProgram.kernel;


		int work_dim = 3;
				
		size_t local_work_size[3];
		local_work_size[3];
		local_work_size[0] = threadsPerTile;
		local_work_size[1] = 1;
		local_work_size[2] = 1;
		
		size_t global_work_size[3];
		global_work_size[0] = threadsPerTile;
		global_work_size[1] = numTiles[0];
		global_work_size[2] = numTiles[1];
				
		clb_check_error_simple(clSetKernelArg(binningKernel, 0, sizeof(cl_mem), &per_block_image2));
		clb_check_error_simple(clSetKernelArg(binningKernel, 1, 2 * sizeof(int), &numTiles));
		clb_check_error_simple(clSetKernelArg(binningKernel, 2, 2 * sizeof(int), &num_blocks));
		clb_check_error_simple(clSetKernelArg(binningKernel, 3, sizeof(cl_mem), &lights));
		clb_check_error_simple(clSetKernelArg(binningKernel, 4, sizeof(int), &numLights));
		clb_check_error_simple(clSetKernelArg(binningKernel, 5, sizeof(cl_mem), &lightBins));
		clb_check_error_simple(clSetKernelArg(binningKernel, 6, sizeof(cl_mem), &binCounts));
		clb_check_error_simple(clSetKernelArg(binningKernel, 7, sizeof(float), &zNear));
		clb_check_error_simple(clSetKernelArg(binningKernel, 8, sizeof(float), &Q1));
		clb_check_error_simple(clSetKernelArg(binningKernel, 9, sizeof(float), &Q2));
		clb_check_error_simple(clSetKernelArg(binningKernel, 10, sizeof(float), &C));
		clb_check_error_simple(clSetKernelArg(binningKernel, 11, sizeof(float), &D));
		clb_check_error_simple(clSetKernelArg(binningKernel, 12, sizeof(cl_sampler), &nearestSampler));
		clb_check_error_simple(clSetKernelArg(binningKernel, 13, 2 * sizeof(int), &imageSize));

		error = clEnqueueNDRangeKernel(queue, binningKernel, work_dim, NULL,\
										global_work_size, local_work_size, 0, NULL, &kernelFinish2Event);
		clb_check_error_simple(error);
		
		
	}
	
#if 1
	{
		
		cl_kernel kernel = computeProgram.kernel;
		float specularPowerScale = 300.0f;
		
		size_t local_work_size[work_dim];
		local_work_size[0] = block_size[0];
		local_work_size[1] = block_size[1];
		size_t global_work_size[work_dim];
		global_work_size[0] = num_blocks[0] * block_size[0];
		global_work_size[1] = num_blocks[1] * block_size[1];
		
		//printf("per pixel: global work size = %d, %d.  Local work size = %d, %d\n",
		//	   global_work_size[0], global_work_size[1], local_work_size[0], local_work_size[1]);
		
		int binOffset = numLights;
						
		clb_check_error_simple(clSetKernelArg(kernel, 0, sizeof(cl_mem), &normSpecImage));
		clb_check_error_simple(clSetKernelArg(kernel, 1, sizeof(cl_mem), &depthImage));
		clb_check_error_simple(clSetKernelArg(kernel, 2, sizeof(cl_mem), &diffuseImage));
		clb_check_error_simple(clSetKernelArg(kernel, 3, sizeof(cl_mem), &specularImage));
		clb_check_error_simple(clSetKernelArg(kernel, 4, sizeof(float), &specularPowerScale));
		clb_check_error_simple(clSetKernelArg(kernel, 5, sizeof(float), &Q1));
		clb_check_error_simple(clSetKernelArg(kernel, 6, sizeof(float), &Q2));
		clb_check_error_simple(clSetKernelArg(kernel, 7, sizeof(float), &C));
		clb_check_error_simple(clSetKernelArg(kernel, 8, sizeof(float), &D));

		clb_check_error_simple(clSetKernelArg(kernel, 9, sizeof(cl_sampler), &nearestSampler));
		clb_check_error_simple(clSetKernelArg(kernel, 10, 2 * sizeof(int), &imageSize));
		clb_check_error_simple(clSetKernelArg(kernel, 11, sizeof(float), &rot));
		clb_check_error_simple(clSetKernelArg(kernel, 12, sizeof(cl_mem), &per_block_image1));
		clb_check_error_simple(clSetKernelArg(kernel, 13, sizeof(cl_mem), &per_block_image2));
		
				
		clb_check_error_simple(clSetKernelArg(kernel, 14, sizeof(cl_mem), &lightBins));
		clb_check_error_simple(clSetKernelArg(kernel, 15, sizeof(cl_mem), &binCounts));
		clb_check_error_simple(clSetKernelArg(kernel, 16, sizeof(cl_mem), &lights));
		clb_check_error_simple(clSetKernelArg(kernel, 17, sizeof(int), &binOffset));
		clb_check_error_simple(clSetKernelArg(kernel, 18, sizeof(int), &numTiles[0]));


		error = clEnqueueNDRangeKernel(queue, kernel, work_dim, NULL, global_work_size, local_work_size, 0, NULL, &kernelFinish3Event);
		clb_check_error_simple(error);
		
	}
#endif
	
   // release the objects back to OpenGL
   error |= clEnqueueReleaseGLObjects(queue, 1, &diffuseImage, 0, 0, 0);
   error |= clEnqueueReleaseGLObjects(queue, 1, &specularImage, 0, 0, 0);
   error |= clEnqueueReleaseGLObjects(queue, 1, &normSpecImage, 0, 0, 0);
   error |= clEnqueueReleaseGLObjects(queue, 1, &depthImage, 0, 0, 0);							   
	
	clFinish(queue);
	clb_check_error_simple(error);

	clb_print_event_profiling_info(kernelFinish1Event, "block kernel");
	clb_print_event_profiling_info(kernelFinish2Event, "bin kernel");
	clb_print_event_profiling_info(kernelFinish3Event, "pixel kernel");

}

-(void)generateLightBuffer {
		
	
	glBindFramebuffer(GL_FRAMEBUFFER, fbo);
	rb_check_error_simple();

	glFramebufferTexture2D(GL_FRAMEBUFFER, GL_COLOR_ATTACHMENT0, GL_TEXTURE_2D, lightbufferDiffuseTextureID, 0);
	glFramebufferTexture2D(GL_FRAMEBUFFER, GL_COLOR_ATTACHMENT1, GL_TEXTURE_2D, lightbufferSpecularTextureID, 0);
	glFramebufferTexture2D(GL_FRAMEBUFFER, GL_DEPTH_ATTACHMENT, GL_TEXTURE_2D, 0, 0);
	
	rb_check_error_simple();
	
	GLenum buffers[2] = { GL_COLOR_ATTACHMENT0, GL_COLOR_ATTACHMENT1};
	glDrawBuffers(2, buffers);
	
	rb_check_error_simple();
	
	GLuint status = glCheckFramebufferStatus(GL_FRAMEBUFFER);
	if ( status != GL_FRAMEBUFFER_COMPLETE ) {
		NSLog(@"Incomplete lightbuffer framebuffer %s", rbunny_framebuffer_status_to_string(status) );
	}				

	[simpleGenLightbufferShader useProgram];
	glUniform1i([simpleGenLightbufferShader locationForUniform:@"normalsSpecSampler"], 0);
	glUniform1i([simpleGenLightbufferShader locationForUniform:@"depthSampler"], 1);

	glUniform1f([simpleGenLightbufferShader locationForUniform:@"zNear"], zNear);
	glUniform1f([simpleGenLightbufferShader locationForUniform:@"zFar"], zFar);
	
	GLfloat imageSize[2];
	imageSize[0] = viewportWidth;
	imageSize[1] = viewportHeight;
	glUniform2fv([simpleGenLightbufferShader locationForUniform:@"imageSize"], 1, imageSize);

	glActiveTexture(GL_TEXTURE0);
	glBindTexture(GL_TEXTURE_2D, normalTextureID);
	
	glActiveTexture(GL_TEXTURE1);
	glBindTexture(GL_TEXTURE_2D, depthTextureID);

	[self drawSquare: [simpleGenLightbufferShader getLocationForAttribute:@"position_in"]];
	
	glActiveTexture(GL_TEXTURE1);
	glBindTexture(GL_TEXTURE_2D, 0);

	glActiveTexture(GL_TEXTURE0);
	glBindTexture(GL_TEXTURE_2D, 0);

	glFramebufferTexture2D(GL_FRAMEBUFFER, GL_COLOR_ATTACHMENT0, GL_TEXTURE_2D, 0, 0);
	glFramebufferTexture2D(GL_FRAMEBUFFER, GL_COLOR_ATTACHMENT1, GL_TEXTURE_2D, 0, 0);
	glFramebufferTexture2D(GL_FRAMEBUFFER, GL_DEPTH_ATTACHMENT, GL_TEXTURE_2D, 0, 0);
	
	glDrawBuffer(GL_COLOR_ATTACHMENT0);

	
}

-(void)depthBlit {

	glBindFramebuffer(GL_FRAMEBUFFER, fbo);
	
	glDisable(GL_DEPTH_TEST);	
	glDepthMask(GL_FALSE);

	glFramebufferTexture2D(GL_FRAMEBUFFER, GL_COLOR_ATTACHMENT0, GL_TEXTURE_2D, depthTexturePrimeID, 0);
	
	glActiveTexture(GL_TEXTURE0);
	glBindTexture(GL_TEXTURE_2D, depthTextureID);

	[blitShader useProgram];
	
	int i = [blitShader getLocationForAttribute:@"position_in"];
	glUniform1i([blitShader locationForUniform:@"sampler"], 0);	
	glUniform1i([blitShader locationForUniform:@"alphaOnly"], 0);	

	[self drawSquare: i];
		
	glBindTexture(GL_TEXTURE_2D, 0);
	
	glFramebufferTexture2D(GL_FRAMEBUFFER, GL_COLOR_ATTACHMENT0, GL_TEXTURE_2D, 0, 0);
	glBindFramebuffer(GL_FRAMEBUFFER, 0);

	glEnable(GL_DEPTH_TEST);
	glDepthMask(GL_TRUE);

	rb_check_error_simple();
	
	
}

-(void)keyDown:(NSEvent *)theEvent {

	
	NSString *characters = [theEvent characters];
	if ([characters length]) {
				
		switch ( [characters characterAtIndex:0] ) {
			case 'z':
				finalBlitMode++;
				finalBlitMode %= kNumFinalBlitModes;
				break;
			case 'x':
				finalBlitMode--;
				if ( finalBlitMode < 0 ) finalBlitMode += kNumFinalBlitModes;
				break;
				
		}
		
	}
	
}

-(void)finalBlit {
	
	glBindFramebuffer(GL_FRAMEBUFFER, 0);
	
	glDisable(GL_DEPTH_TEST);
	glDepthMask(GL_FALSE);
	//glClear(GL_COLOR_BUFFER_BIT | GL_DEPTH_BUFFER_BIT);
	
	glActiveTexture(GL_TEXTURE0);
	
	char *modeString;
	GLint blitTexture;
	switch( finalBlitMode ) {
		case kFinalBlitModeColor: 
			blitTexture = colorTextureID;
			modeString = "Color";
			break;
		case kFinalBlitModeNormal:
			blitTexture = normalTextureID;
			modeString = "Normal";
			break;
		case kFinalBlitModeDepth: 
			blitTexture = depthTextureID;
			modeString = "Depth";
			break;
		case kFinalBlitModeDiffuse: 
			blitTexture = lightbufferDiffuseTextureID;
			modeString = "Diffuse";
			break;
		case kFinalBlitModeSpecular: 
			blitTexture = lightbufferSpecularTextureID;
			modeString = "Specular";
			break;
		case kFinalBlitModeDepthPrime:
			blitTexture = depthTexturePrimeID;
			modeString = "Depth Prime";
			break;
		case kFinalBlitModeDebug1:
			blitTexture = lightbufferDiffuseTextureID;
			modeString = "Debug 1";
			break;
		case kFinalBlitModeDebug2:
			blitTexture = lightbufferSpecularTextureID;
			modeString = "Debug 2";
			break;			
		default: 
			blitTexture = 0; 
			modeString = "???";
			break;
	}	
	
	glBindTexture(GL_TEXTURE_2D, blitTexture);
	
	[blitShader useProgram];
	
	BOOL isDebug = (finalBlitMode == kFinalBlitModeDebug2) || (finalBlitMode == kFinalBlitModeDebug1);
	
	int i = [blitShader getLocationForAttribute:@"position_in"];
	glUniform1i([blitShader locationForUniform:@"sampler"], 0);	
	glUniform1i([blitShader locationForUniform:@"alphaOnly"], isDebug);	

	[self drawSquare: i];
	
	glBindTexture(GL_TEXTURE_2D, 0);
	
	[self drawHUD];
	
	glEnable(GL_DEPTH_TEST);
	glDepthMask(GL_TRUE);
	
	rb_check_error_simple();
		
	
}

-(void)drawRect:(NSRect)dirtyRect {
			
	[self.openGLContext makeCurrentContext];
	
	/* compute time delta */
	struct timeval currentTime;
	gettimeofday(&currentTime, NULL);
	double timeDiff = (double)( currentTime.tv_usec - lastTime.tv_usec ) * (double)(1.0 / 1000000.0) + (double)( currentTime.tv_sec - lastTime.tv_sec );
	if ( timeDiff < 0.0 ) timeDiff = 0.0;
	if ( timeDiff > 1.0 ) timeDiff = 1.0;
	lastTime = currentTime;
	/* ------------------- */
	
	rot += 0.2 * timeDiff;
				
	fovyInDegrees = 35.0f;
	aspect = (float)viewportWidth / (float)viewportHeight;
	zNear = 0.1f;
	zFar  = 400.0f;
	top = zNear * tanf(fovyInDegrees * M_PI / 360.0);
	right = top * aspect;
    bottom = -top;
    left = -top * aspect;
		
	mat4 modelview = [self getModelview];
	
	[self drawPass: kPassGBuffer];	// about 4ms
	[self depthBlit];				// about 3ms
	
	glFinish();
	
	[particleSystem stepParticleSystem:  0.05f * timeDiff reset: NO modelviewMatrix: modelview];	
	
	[self generateLightBufferOpenCL]; // generate the light buffer
	
	clFinish(simpleCompute.queue);

	[self drawPass: kPassShade];	// generate the color texture (6ms)
	[self finalBlit];				// blit the color texture to the screen and do any final effects (2ms)
	
	[self.openGLContext flushBuffer];
	
	count++;

}

- (void)drawHUD {

	[basicShader useProgram];
	
	glEnable(GL_BLEND);
	glDisable(GL_DEPTH_TEST);
	glDepthMask(GL_FALSE);
	glBlendFunc(GL_ONE, GL_ONE_MINUS_SRC_ALPHA);
	
	mat4 projection = rb_mat4_ortho(0, self.bounds.size.width, self.bounds.size.height, 0, -10, 10);
	mat4 modelview  = rb_mat4_translate(make_vec3(10, 10, 0));
	
	mat4 modelviewProjection = rb_mat4_mul(projection, modelview);
	
	glUniformMatrix4fv( [basicShader locationForUniform:@"modelViewProjectionMatrix"], 1, 0, modelviewProjection.entries );
	glUniform1i([basicShader locationForUniform:@"sampler"], 0);
	glUniform4f([basicShader locationForUniform:@"alpha"], 1.0f, 1.0f, 1.0f, 1.0f);
	
	rb_check_error_simple();
	
	[GLFont mapVerticesToAttributeLocation: [basicShader getLocationForAttribute:@"position_in"] ];
	[GLFont mapTexCoordsToAttributeLocation: [basicShader getLocationForAttribute:@"uv_in"] ];	
	
	rb_check_error_simple();

	[font drawString:[NSString stringWithFormat:@"Frames per second: %d", fps]];
	
	rb_check_error_simple();
	
	glDisable(GL_BLEND);
	
	/*glActiveTexture(GL_TEXTURE0);
	glBindTexture(GL_TEXTURE_2D, normalTextureID);
	
	[self visualizeColorWithImageSize: CGSizeMake(viewportWidth, viewportHeight) dstRect: CGRectMake(4, 4, 128, 1.0f / aspect * 128)];
	
	glBindTexture(GL_TEXTURE_2D, depthTextureID);
	[self visualizeColorWithImageSize: CGSizeMake(viewportWidth, viewportHeight) dstRect: CGRectMake(128 + 8, 4, 128, 1.0f / aspect * 128)];
	
	glBindTexture(GL_TEXTURE_2D, lightbufferDiffuseTextureID);
	[self visualizeColorWithImageSize: CGSizeMake(viewportWidth, viewportHeight) dstRect: CGRectMake(256 + 32, 4, 128, 1.0f / aspect * 128)];

	glBindTexture(GL_TEXTURE_2D, lightbufferSpecularTextureID);
	[self visualizeColorWithImageSize: CGSizeMake(viewportWidth, viewportHeight) dstRect: CGRectMake(384 + 36, 4, 128, 1.0f / aspect * 128)];

	glBindTexture(GL_TEXTURE_2D, colorTextureID);
	[self visualizeColorWithImageSize: CGSizeMake(viewportWidth, viewportHeight) dstRect: CGRectMake(512 + 60, 4, 128, 1.0f / aspect * 128)];

	glBindTexture(GL_TEXTURE_2D, depthTexturePrimeID);
	[self visualizeColorWithImageSize: CGSizeMake(viewportWidth, viewportHeight) dstRect: CGRectMake(640 + 64, 4, 128, 1.0f / aspect * 128)];
*/
	
	glDepthMask(GL_TRUE);
	glEnable(GL_DEPTH_TEST);
	
}

-(void)mouseUp:(NSEvent *)theEvent {

	
}

- (void)rebuildFramebufferObjects {
	
	glFinish();
	clFinish(simpleCompute.queue);
	
	GLint imageSize[2];
	imageSize[0] = viewportWidth;
	imageSize[1] = viewportHeight;	
	num_blocks[0] = ( imageSize[0] + block_size[0] - 1 ) / block_size[0];
	num_blocks[1] = ( imageSize[1] + block_size[1] - 1 ) / block_size[1];
		
	if ( colorTextureID != 0 ) {
		glDeleteTextures(1, &colorTextureID);
	}
	if ( normalTextureID != 0 ) {
		glDeleteTextures(1, &normalTextureID);
	}
	if ( depthTextureID != 0 ) {
		glDeleteTextures(1, &depthTextureID);
	}
	if ( depthTexturePrimeID != 0 ) {
		glDeleteTextures(1, &depthTexturePrimeID);
	}
	
	
	if ( lightbufferDiffuseTextureID != 0 ) {
		glDeleteTextures(1, &lightbufferDiffuseTextureID);
	}
	if ( lightbufferSpecularTextureID != 0 ) {
		glDeleteTextures(1, &lightbufferSpecularTextureID);
	}
	
	
	
	if ( fbo == 0 ) {
		glGenFramebuffers(1, &fbo);
	}
		
	glBindFramebuffer(GL_FRAMEBUFFER, fbo);
	
	glFramebufferTexture2D(GL_FRAMEBUFFER, GL_COLOR_ATTACHMENT0, GL_TEXTURE_2D, 0, 0);
	glFramebufferTexture2D(GL_FRAMEBUFFER, GL_DEPTH_ATTACHMENT, GL_TEXTURE_2D, 0, 0);

	
	rb_check_error_simple();
	
	glActiveTexture(GL_TEXTURE0);

	
	glGenTextures(1, &depthTextureID);

	rb_check_error_simple();
	
	glBindTexture(GL_TEXTURE_2D, depthTextureID);

	rb_check_error_simple();

	
	glTexImage2D(GL_TEXTURE_2D, 0, GL_DEPTH_COMPONENT24, viewportWidth, viewportHeight, 0, GL_DEPTH_COMPONENT, GL_UNSIGNED_INT, NULL);
	
	rb_check_error_simple();
	
	glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_WRAP_S, GL_CLAMP);
	glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_WRAP_T, GL_CLAMP);
	glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_MIN_FILTER, GL_NEAREST);
	glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_MAG_FILTER, GL_NEAREST);
	
	rb_check_error_simple();

	glFramebufferTexture2D(GL_FRAMEBUFFER, GL_DEPTH_ATTACHMENT, GL_TEXTURE_2D, depthTextureID, 0);
	
	rb_check_error_simple();
	
	// generate color texture
	glGenTextures(1, &colorTextureID);
	glBindTexture(GL_TEXTURE_2D, colorTextureID);
	glTexImage2D(GL_TEXTURE_2D, 0, GL_RGBA8, viewportWidth, viewportHeight, 0, GL_RGBA, GL_UNSIGNED_BYTE, NULL);
	glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_WRAP_S, GL_CLAMP);
	glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_WRAP_T, GL_CLAMP);
	glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_MIN_FILTER, GL_NEAREST);
	glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_MAG_FILTER, GL_NEAREST);
	rb_check_error_simple();
	glFramebufferTexture2D(GL_FRAMEBUFFER, GL_COLOR_ATTACHMENT0, GL_TEXTURE_2D, colorTextureID, 0);
	rb_check_error_simple();
	
	GLuint status = glCheckFramebufferStatus(GL_FRAMEBUFFER);
	if ( status != GL_FRAMEBUFFER_COMPLETE ) {
		NSLog(@"Incomplete framebuffer (color) %s", rbunny_framebuffer_status_to_string(status) );
	}
	
	rb_check_error_simple();

	// generate normal texture
	glGenTextures(1, &normalTextureID);
	glBindTexture(GL_TEXTURE_2D, normalTextureID);
	
	glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_WRAP_S, GL_CLAMP);
	glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_WRAP_T, GL_CLAMP);
	glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_MIN_FILTER, GL_NEAREST);
	glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_MAG_FILTER, GL_NEAREST);	
	
	// normals cannot be RGBA888
	glTexImage2D(GL_TEXTURE_2D, 0, GL_RGBA, viewportWidth, viewportHeight, 0, GL_RGBA, GL_UNSIGNED_BYTE, NULL);
	
	rb_check_error_simple();

	rb_check_error_simple();
	glFramebufferTexture2D(GL_FRAMEBUFFER, GL_COLOR_ATTACHMENT0, GL_TEXTURE_2D, normalTextureID, 0);
	rb_check_error_simple();
	
	status = glCheckFramebufferStatus(GL_FRAMEBUFFER);
	if ( status != GL_FRAMEBUFFER_COMPLETE ) {
		NSLog(@"Incomplete framebuffer (normal) %s", rbunny_framebuffer_status_to_string(status) );
	}
		
	// generate light buffer textures
	glGenTextures(1, &lightbufferDiffuseTextureID);
	rb_check_error_simple();
	
	glBindTexture(GL_TEXTURE_2D, lightbufferDiffuseTextureID);
	glTexImage2D(GL_TEXTURE_2D, 0, GL_RGBA16F_ARB, viewportWidth, viewportHeight, 0, GL_RGBA, GL_HALF_APPLE, NULL);
	glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_WRAP_S, GL_CLAMP);
	glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_WRAP_T, GL_CLAMP);
	glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_MIN_FILTER, GL_NEAREST);
	glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_MAG_FILTER, GL_NEAREST);
	rb_check_error_simple();
	
	glGenTextures(1, &lightbufferSpecularTextureID);
	rb_check_error_simple();

	glBindTexture(GL_TEXTURE_2D, lightbufferSpecularTextureID);
	glTexImage2D(GL_TEXTURE_2D, 0, GL_RGBA16F_ARB, viewportWidth, viewportHeight, 0, GL_RGBA, GL_HALF_APPLE, NULL);
	glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_WRAP_S, GL_CLAMP);
	glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_WRAP_T, GL_CLAMP);
	glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_MIN_FILTER, GL_NEAREST);
	glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_MAG_FILTER, GL_NEAREST);
	rb_check_error_simple();
	
	
	// generate normal texture
	glGenTextures(1, &depthTexturePrimeID);
	glBindTexture(GL_TEXTURE_2D, depthTexturePrimeID);
	glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_WRAP_S, GL_CLAMP);
	glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_WRAP_T, GL_CLAMP);
	glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_MIN_FILTER, GL_NEAREST);
	glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_MAG_FILTER, GL_NEAREST);	
	glTexImage2D(GL_TEXTURE_2D, 0, GL_RGBA32F_ARB, viewportWidth, viewportHeight, 0, GL_RGBA, GL_FLOAT, NULL);
	
	rb_check_error_simple();
	
	glBindFramebuffer(GL_FRAMEBUFFER, 0);
	glFinish();
	
	clFinish(simpleCompute.queue);

	cl_int error = CL_SUCCESS;
	if (normSpecImage)	error |= clEnqueueAcquireGLObjects(simpleCompute.queue, 1, &normSpecImage, 0, 0, 0);
	if (depthImage)		error |= clEnqueueAcquireGLObjects(simpleCompute.queue, 1, &depthImage, 0, 0, 0);
	if (diffuseImage)	error |= clEnqueueAcquireGLObjects(simpleCompute.queue, 1, &diffuseImage, 0, 0, 0);
	if (specularImage)	error |= clEnqueueAcquireGLObjects(simpleCompute.queue, 1, &specularImage, 0, 0, 0);
	clb_check_error_simple(error);
	
	if ( diffuseImage )		clReleaseMemObject(diffuseImage);
	if ( specularImage )	clReleaseMemObject(specularImage);
	if ( normSpecImage )	clReleaseMemObject(normSpecImage);
	if ( depthImage )		clReleaseMemObject(depthImage);		
	
	diffuseImage		= [self createComputeResultFromTexture: lightbufferDiffuseTextureID		mode:	CL_MEM_WRITE_ONLY ];
	specularImage		= [self createComputeResultFromTexture: lightbufferSpecularTextureID	mode:	CL_MEM_WRITE_ONLY ];
	normSpecImage		= [self createComputeResultFromTexture: normalTextureID					mode:	CL_MEM_READ_ONLY ];
	depthImage			= [self createComputeResultFromTexture: depthTexturePrimeID				mode:	CL_MEM_READ_ONLY ];
		
	if ( per_block_image1 ) error |= clReleaseMemObject(per_block_image1);
	if ( per_block_image2 ) error |= clReleaseMemObject(per_block_image2);
	clb_check_error_simple(error);
	
	cl_image_format format1;
	format1.image_channel_order = CL_RGBA;
	format1.image_channel_data_type = CL_FLOAT;

	cl_image_format format2;
	format2.image_channel_order = CL_RGBA;
	format2.image_channel_data_type = CL_FLOAT;
	
	per_block_image1	= [self createImageForOpenCLOnlyWithSize:CGSizeMake(num_blocks[0] , num_blocks[1]) format: format1 ];
	per_block_image2	= [self createImageForOpenCLOnlyWithSize:CGSizeMake(num_blocks[0] , num_blocks[1]) format: format2 ];
	
	error = clEnqueueReleaseGLObjects(simpleCompute.queue, 1, &normSpecImage, 0, 0, 0);
	error |= clEnqueueReleaseGLObjects(simpleCompute.queue, 1, &depthImage, 0, 0, 0);
	error |= clEnqueueReleaseGLObjects(simpleCompute.queue, 1, &diffuseImage, 0, 0, 0);
	error |= clEnqueueReleaseGLObjects(simpleCompute.queue, 1, &specularImage, 0, 0, 0);
	clb_check_error_simple(error);

	
}

- (void)prepareOpenGL {
		
	NSLog(@"Prepare OpenGL");
	
	[self.openGLContext makeCurrentContext];
	
	[self loadOpenCL];
	[self loadShaders];
		
	model1 = [[IndexedModel alloc] initWithFile: [[NSBundle mainBundle] pathForResource:@"volcano5" ofType:@"obj"] ];
	
	[model1 mapBuffer: IMODEL_BUFFER_VERTEX			toAttributeLocation: [tangentShader getLocationForAttribute:@"position"]];
	[model1 mapBuffer: IMODEL_BUFFER_UV				toAttributeLocation: [tangentShader getLocationForAttribute:@"uv"]];
	[model1 mapBuffer: IMODEL_BUFFER_NORMAL			toAttributeLocation: [tangentShader getLocationForAttribute:@"normal"]];
	[model1 mapBuffer: IMODEL_BUFFER_TANGENT		toAttributeLocation: [tangentShader getLocationForAttribute:@"tangent"]];
	[model1 createVertexArray];
	
	model2 = [[IndexedModel alloc] initWithFile: [[NSBundle mainBundle] pathForResource:@"volcano5" ofType:@"obj"] ];
	[model2 mapBuffer: IMODEL_BUFFER_VERTEX			toAttributeLocation: [lightBufferShader getLocationForAttribute:@"position"]];
	[model2 mapBuffer: IMODEL_BUFFER_UV				toAttributeLocation: [lightBufferShader getLocationForAttribute:@"uv"]];
	[model2 createVertexArray];
	
	rb_check_error_simple();
	
	
	texture = [[GLTexture alloc] initWithFile: [[NSBundle mainBundle] pathForResource:@"volcano-texture" ofType:@"png"]
										flags: kGLTextureFilterModeLinear|kGLTextureWrapModeRepeat|kGLTextureGenerateMipMaps];
    	
	normals = [[GLTexture alloc] initWithFile: [[NSBundle mainBundle] pathForResource:@"leather-normals" ofType:@"png"]
										flags: kGLTextureFilterModeLinear|kGLTextureWrapModeRepeat|kGLTextureGenerateMipMaps ];
		
	
	GLubyte whiteColor[4] = { 255, 255, 255, 255 };
	GLubyte blackColor[4] = { 0, 0, 0, 255 };
	
	font = [[GLFont alloc] initWithFontName:[[NSFont systemFontOfSize: 16] fontName] size: 16 color: whiteColor stroke: 1.0f strokeColor: blackColor];
	
	glClearDepth(1.0f);
    glClearColor(0.5f, 0.5f, 0.5f, 1.0f);
	glEnable(GL_CULL_FACE);	
		
	rb_check_error_simple();
		
	viewportWidth = self.bounds.size.width;
	viewportHeight = self.bounds.size.height;
	
	per_block_image1 = 0;
	per_block_image2 = 0;
	colorTextureID = 0;
	depthTextureID = 0;
	depthTexturePrimeID = 0;
	normalTextureID = 0;
	lightbufferDiffuseTextureID = 0;
	lightbufferSpecularTextureID = 0;
	diffuseImage = 0;
	specularImage = 0;
	normSpecImage =0;
	depthImage=0;
	
	[self rebuildFramebufferObjects];

}

-(void)reshape {
	
	[self.openGLContext makeCurrentContext];
	glFinish();

	viewportWidth = self.bounds.size.width;
	viewportHeight = self.bounds.size.height;
	glViewport(0, 0, viewportWidth, viewportHeight);
	
	clFinish(simpleCompute.queue);
		
	[self rebuildFramebufferObjects];
	
	[self setNeedsDisplay: YES];
	
}

@end
